#include "..\..\include\pcval.h"
#include <CL\cl.h>

extern cl_context			g_device_context;
extern cl_command_queue		g_device_command_queue;
extern cl_kernel			g_kernel_list[pcvClsMAX][pcvFuncMax];

const char* pcvAdd_8u_C1Rs_Kernel_Src[] = {
"__kernel void pcvAdd_8u_C1Rs_Kernel(",
"	__global unsigned char*	img_src1,",
"	__global unsigned char*	img_src2,",
"	__global unsigned char*	img_dst,",
"	int host_size_width)",
"{",
"	size_t					x = get_global_id(0);",
"	size_t					y = get_global_id(1);",
"	size_t					offset = y * host_size_width + x;",
"	img_dst[offset] = add_sat(img_src1[offset], img_src2[offset]);",
"}"
};

PCVAPI(PcvStatus, pcvAdd_8u_C1Rs_Build, (
	))
{
	cl_int					errcode = CL_SUCCESS;

	if (g_device_context == NULL ||
		g_kernel_list[pcvClsAL][pcvFuncALAdd_8u_C1Rs] != NULL)
	{
		return pcvStsErr;
	}

	// OpenCL supports runtime code compilation. First we build the OpenCL
	// program from source code
	cl_program program = clCreateProgramWithSource(g_device_context, 11, pcvAdd_8u_C1Rs_Kernel_Src, NULL, &errcode);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	errcode = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (errcode != CL_SUCCESS)
	{
		clReleaseProgram(program);
		return pcvStsErr;
	}

	// Create a handle to the compiled OpenCL function (Kernel)
	g_kernel_list[pcvClsAL][pcvFuncALAdd_8u_C1Rs] = clCreateKernel(program, "pcvAdd_8u_C1Rs_Kernel", &errcode);
	if (errcode != CL_SUCCESS)
	{
		clReleaseProgram(program);
		return pcvStsErr;
	}

	errcode = clReleaseProgram(program);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	return pcvStsNoErr;
}

PCVAPI(PcvStatus, pcvAdd_8u_C1Rs, (
	const void*				device_img_src1,
	const void*				device_img_src2,
	void*					device_img_dst,
	const PcvSize			host_size,
	const PcvRect			host_roi
	))
{
	cl_int					errcode = CL_SUCCESS;

	if (g_device_command_queue == NULL)
	{
		return pcvStsNoInitErr;
	}

	cl_kernel kernel = g_kernel_list[pcvClsAL][pcvFuncALAdd_8u_C1Rs];
	if (kernel == NULL)
	{
		return pcvStsNotBuildErr;
	}

	// In the next step we associate the GPU memory with the Kernel arguments
	int						offset = host_roi.y * host_size.width + host_roi.x;
	unsigned char*			arg = (unsigned char*)device_img_src1 + offset;
	errcode = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&device_img_src1);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	arg = (unsigned char*)device_img_src2 + offset;
	errcode = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&device_img_src2);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	arg = (unsigned char*)device_img_dst + offset;
	errcode = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&device_img_dst);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	errcode = clSetKernelArg(kernel, 3, sizeof(int), (void*)&(host_size.width));
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	// Then we launch the Kernel on the GPU
	size_t					work_size[2] = {host_size.width, host_size.height};
	cl_event				event_ret;
	errcode = clEnqueueNDRangeKernel(g_device_command_queue, kernel, 2, NULL, work_size, NULL, 0, NULL, &event_ret);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	errcode = clWaitForEvents(1, &event_ret);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	cl_event_info			event_info = CL_EVENT_COMMAND_EXECUTION_STATUS;
	cl_int					event_status;
	errcode = clGetEventInfo(event_ret, event_info, sizeof(cl_int), &event_status, NULL);
	if (errcode != CL_SUCCESS || event_status != CL_COMPLETE)
	{
		return pcvStsErr;
	}

	return pcvStsNoErr;
}
